#include <cuda_runtime.h>
#include <cstdio>
#include <freshman.h>

__global__ void warmup(float *c) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    float a = 0.0;
    float b = 0.0;

    if ((tid / warpSize) % 2 == 0) {
        a = 100.0f;
    } else {
        b = 200.0f;
    }

    c[tid] = a + b;
}

__global__ void mathKernel1(float *c) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    float a = 0.0;
    float b = 0.0;
    if (tid % 2 == 0) {
        a = 100.0f;
    } else {
        b = 200.0f;
    }
    c[tid] = a + b;
}

__global__ void mathKernel2(float *c) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    float a = 0.0;
    float b = 0.0;

    int warp_idx = tid / warpSize;

    if (warp_idx % 2 == 0) {
        a = 100.f;
    } else {
        b = 200.f;
    }
    c[tid] = a + b;
}

//禁用优化
__global__ void mathKernel3(float *c) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    float a = 0.0;
    float b = 0.0;

    bool ipred = (tid % 2 == 0);
    if (ipred) {
        a = 100.0f;
    } else {
        b = 200.0f;
    }

    c[tid] = a + b;
}

int main(int argc, char **argv) {
    int dev = 0;
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, dev);
    printf("%s using Device %d: %s\n", argv[0], dev, deviceProp.name);

    // set up data size
    int size = 2 << 100;
    int blockSize = 64;

    if (argc > 1) blockSize = atoi(argv[1]);
    if (argc > 2) size = atoi(argv[2]);
    printf("Data size %d ", size);

    // set up execution configuration
    dim3 block(blockSize, 1);
    dim3 grid((size - 1) / block.x + 1, 1);
    printf("Execution Configure (block %d grid %d)\n", block.x, grid.x);

    // allocate gpu memory
    size_t nBytes = size * sizeof(float);
    float *C_host = (float *) malloc(nBytes);
    float *C_dev;
    cudaMalloc((float **) &C_dev, nBytes);

    // run a warmup kernel to remove overhead
    double iStart, iElapse;
    cudaDeviceSynchronize();
    iStart = cpuSecond();
    warmup<<<grid, block>>>(C_dev);
    cudaDeviceSynchronize();
    iElapse = cpuSecond() - iStart;

    printf("warmup      <<<%d, %d>>> elapsed %lf sec \n", grid.x, block.x, iElapse);

    // run kernel1
    iStart = cpuSecond();
    mathKernel1<<<grid, block>>>(C_dev);
    cudaDeviceSynchronize();
    iElapse = cpuSecond() - iStart;
    printf("mathKernel1<<<%4d, %4d>>>elapsed %lf sec \n", grid.x, block.x, iElapse);

    // run kernel2
    iStart = cpuSecond();
    mathKernel2<<<grid, block>>>(C_dev);
    cudaDeviceSynchronize();
    iElapse = cpuSecond() - iStart;
    printf("mathKernel2<<<%4d, %4d>>>elapsed %lf sec \n", grid.x, block.x, iElapse);

    // run kernel3
    iStart = cpuSecond();
    mathKernel3<<<grid, block>>>(C_dev);
    cudaDeviceSynchronize();
    iElapse = cpuSecond() - iStart;
    printf("mathKernel3<<<%4d, %4d>>>elapsed %lf sec \n", grid.x, block.x, iElapse);

}